/*
 * Copyright (c) 2020-2022, NVIDIA CORPORATION. All rights reserved.
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 *     http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

// Original from https://github.com/NVIDIA/TensorRT/commit/09dce45f919be40d27b2b650e716debb3740d410
// Modified based on FasterTransformer 5.2
#pragma once

#if defined(_MSC_VER)
#pragma warning(push)
// Global initializer accesses extern object cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len
#pragma warning(disable : 26427)
#endif
#include "contrib_ops/cuda/bert/tensorrt_fused_multihead_attention/fused_multihead_attention.h"
#include <cstdint>

namespace onnxruntime {
namespace contrib {
namespace cuda {
struct Fused_multihead_attention_params_v2 {
  // The QKV matrices.
  void* qkv_ptr{};
  // The mask to implement drop-out.
  void* packed_mask_ptr{};
  // The O matrix (output).
  void* o_ptr{};

  // The stride between rows of the Q, K and V matrices.
  int64_t qkv_stride_in_bytes{};
  // The stride between matrices of packed mask.
  int64_t packed_mask_stride_in_bytes{};
  // The stride between rows of O.
  int64_t o_stride_in_bytes;

  // The dimensions.
  int32_t b{};
  int32_t h{};
  int32_t s{};
  int32_t d{};
  // The scaling factors for the kernel.
  uint32_t scale_bmm1{};
  uint32_t scale_softmax{};
  uint32_t scale_bmm2{};

  // Do we use Niall's trick to avoid I2F/F2I in the INT8 kernel.
  // See https://confluence.nvidia.com/pages/viewpage.action?pageId=302779721 for details.
  bool enable_i2f_trick{};

  // array of length b+1 holding prefix sum of actual sequence lengths
  int32_t* cu_seqlens{};

  // use C/32 Format.
  bool interleaved{};
  bool ignore_b1opt{};
  bool force_unroll{};
  bool use_int8_scale_max{};

  // Additional parameters for relative position bias
  // void* packed_relative_position_bias_ptr{};
  int window_num{};
  // int actual_seqlen{};

  void clear() {
    *this = Fused_multihead_attention_params_v2();
  }
};

////////////////////////////////////////////////////////////////////////////////////////////////////
extern unsigned char cubin_fmha_v2_fp16_384_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_32_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_64_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_32_sm75_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_384_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_256_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_192_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_128_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_96_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_64_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_32_64_sm75_cu_cubin[];

extern unsigned char fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin[];
extern unsigned char fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin[];

// NOTE: flash attention kernels for GPT (causal mask)
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin[];

extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin[];

extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin[];

extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin[];

extern unsigned char cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin[];
extern unsigned char cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin[];

extern unsigned int cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_32_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_32_sm75_cu_cubin_len;

extern unsigned int cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_96_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_64_64_sm75_cu_cubin_len;
extern unsigned int cubin_fmha_v2_fp16_32_64_sm75_cu_cubin_len;

extern unsigned int fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len;
extern unsigned int fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len;

extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len;

extern uint32_t cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len;
extern uint32_t cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len;

static const struct FusedMultiHeadAttentionKernelMetaInfoV2 {
  Data_type mDataType;
  unsigned int mS;
  unsigned int mD;
  unsigned int mSM;
  const unsigned char* mCubin;
  unsigned int mCubinSize;
  const char* mFuncName;
  unsigned int mSharedMemBytes;
  unsigned int mThreadsPerCTA;
  unsigned int mUnrollStep;
  bool mInterleaved;
  bool mWithRelativePositionBias = false;
  bool mFlashAttention = false;
  bool mCausalMask = false;
} sMhaKernelMetaInfosV2[] = {
    // Volta
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_64_64_kernel_sm70",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_64_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_64_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_96_64_kernel_sm70",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_96_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_96_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_128_64_kernel_sm70",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_128_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_128_64_kernel_sm70_noloop",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_256_64_kernel_sm70",
     69632,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_256_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_256_64_kernel_sm70_noloop",
     69632,
     256,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_384_64_kernel_sm70",
     69632,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_70,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin,
     fused_multihead_attention_v2_fp16_384_64_kernel_sm70_cubin_len,
     "fused_multihead_attention_v2_fp16_384_64_kernel_sm70_noloop",
     69632,
     256,
     32,
     false,
     false},

    // Turing
    {DATA_TYPE_FP16,
     32,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_32_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_96_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm75_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm75_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm75_kernel",
     28672,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm75_kernel_nl",
     28672,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm75_kernel",
     36864,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm75_kernel_nl",
     36864,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm75_kernel",
     53248,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm75_kernel_nl",
     53248,
     256,
     32,
     false,
     false},

#if CUDA_VERSION >= 11000
    // Ampere
    {DATA_TYPE_FP16,
     32,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_32_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm75_kernel",
     6144,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_96_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm75_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm75_kernel",
     14336,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm75_kernel_nl",
     14336,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm75_kernel",
     18432,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm75_kernel_nl",
     18432,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm75_kernel",
     28672,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm75_kernel_nl",
     28672,
     256,
     32,
     false,
     false},

    {DATA_TYPE_FP16,
     32,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm80_kernel",
     49152,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel_nl",
     32768,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel",
     65536,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel_nl",
     57344,
     256,
     32,
     false,
     false},

    {DATA_TYPE_FP16,
     32,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm80_kernel",
     8192,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm80_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel_nl",
     16384,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel",
     32768,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel_nl",
     32768,
     256,
     32,
     false,
     false},

    // GA10x: sm86 uses sm80 kernels
    {DATA_TYPE_FP16,
     32,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_32_sm80_kernel",
     8192,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_32_sm80_kernel",
     24576,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_32_sm80_kernel_nl",
     16384,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel",
     20480,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_32_sm80_kernel_nl",
     20480,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel",
     32768,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_32_sm80_kernel_nl",
     32768,
     256,
     32,
     false,
     false},

    {DATA_TYPE_FP16,
     32,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_32_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_32_64_sm80_kernel",
     16384,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     96,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_96_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_96_64_sm80_kernel",
     49152,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_128_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     192,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_192_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_192_64_sm80_kernel_nl",
     32768,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel",
     40960,
     128,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     256,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_256_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_256_64_sm80_kernel_nl",
     40960,
     128,
     32,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel",
     65536,
     256,
     0,
     false,
     false},
    {DATA_TYPE_FP16,
     384,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_384_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_384_64_sm80_kernel_nl",
     57344,
     256,
     32,
     false,
     false},
#endif

    // Flash attention
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm70_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm70_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm70_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm70_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm70_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm70_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm70_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm70_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm70_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm70_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm75_kernel_nl",
     12288,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm75_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm75_kernel_nl",
     24576,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm75_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm75_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm70_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_70,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm70_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm70_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm75_kernel",
     65536,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_75,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm75_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm75_kernel_nl",
     65536,
     128,
     64,
     false,
     false,
     true,
     true},

    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm75_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm75_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm75_kernel",
     12288,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm75_kernel_nl",
     10240,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm75_kernel",
     49152,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm75_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm75_kernel",
     24576,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_75,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm75_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm75_kernel_nl",
     20480,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm70_kernel",
     81920,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm70_kernel_nl",
     36864,
     256,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm70_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm70_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm70_kernel",
     36864,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm70_kernel_nl",
     36864,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm70_kernel",
     69632,
     256,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm70_kernel_nl",
     69632,
     256,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm70_kernel",
     20480,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm70_kernel_nl",
     18432,
     128,
     32,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm70_kernel",
     40960,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_70,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm70_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm70_kernel_nl",
     36864,
     128,
     32,
     false,
     false,
     false,
     true},

#if CUDA_VERSION >= 11000
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm80_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm80_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm80_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm80_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm80_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm80_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm80_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm86_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm86_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm86_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm86_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm86_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm86_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm86_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm86_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm89_kernel",
     16384,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     32,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_32_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_32_sm89_kernel_nl",
     16384,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     40,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_40_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_40_sm89_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     64,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_64_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_64_sm89_kernel_nl",
     32768,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm89_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     80,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_80_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_80_sm89_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm89_kernel",
     49152,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     128,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_128_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_128_sm89_kernel_nl",
     49152,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     160,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_160_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_160_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm89_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     256,
     kSM_89,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_256_sm89_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_256_sm89_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm80_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_80,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm80_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm80_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm86_kernel",
     98304,
     128,
     0,
     false,
     false,
     true,
     true},
    {DATA_TYPE_FP16,
     0,
     144,
     kSM_86,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin,
     cubin_fmha_v2_flash_attention_fp16_Causal_0_144_sm86_cu_cubin_len,
     "fmha_v2_flash_attention_fp16_Causal_0_144_sm86_kernel_nl",
     98304,
     128,
     64,
     false,
     false,
     true,
     true},

    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm80_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm80_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm80_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm80_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm80_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm80_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm80_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_80,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm80_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm80_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm86_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm86_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm86_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm86_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm86_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm86_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm86_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_86,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm86_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm86_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_32_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm89_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_64_sm89_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm89_kernel",
     65536,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     128,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_128_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_128_40_sm89_kernel_nl",
     36864,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm89_kernel",
     16384,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     32,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_32_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_32_sm89_kernel_nl",
     12288,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     64,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_64_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_64_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm89_kernel",
     32768,
     128,
     0,
     false,
     false,
     false,
     true},
    {DATA_TYPE_FP16,
     64,
     40,
     kSM_89,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin,
     cubin_fmha_v2_fp16_Causal_64_40_sm89_cu_cubin_len,
     "fmha_v2_fp16_Causal_64_40_sm89_kernel_nl",
     20480,
     128,
     16,
     false,
     false,
     false,
     true}
#endif
};

class FusedMultiHeadAttentionXMMAKernelV2 : public TFusedMultiHeadAttentionXMMAKernel<FusedMultiHeadAttentionKernelMetaInfoV2,
                                                                                      Fused_multihead_attention_params_v2> {
 public:
  FusedMultiHeadAttentionXMMAKernelV2(const FusedMultiHeadAttentionKernelMetaInfoV2* pMetaStart,
                                      unsigned int nMetaCount,
                                      Data_type type,
                                      unsigned int sm) : TFusedMultiHeadAttentionXMMAKernel<FusedMultiHeadAttentionKernelMetaInfoV2,
                                                                                            Fused_multihead_attention_params_v2>(pMetaStart, nMetaCount, type, sm) {
  }

  inline uint64_t hashID(unsigned int s,
                         unsigned int d,
                         bool interleaved,
                         bool unroll,
                         bool with_relative_position_bias,
                         bool flash_attention = false,
                         bool causal_mask = false) const {
    s = flash_attention ? 0 : s;
    return (uint64_t)s << 32 | (uint64_t)d << 8 | (with_relative_position_bias ? 16ull : 0ull) | (flash_attention ? 8ull : 0ull) | (causal_mask ? 4ull : 0ull) | (interleaved ? 2ull : 0ull) | (unroll ? 1ull : 0ull);
  }

  virtual uint64_t hashID(const KernelMeta& kernelMeta) const {
    ORT_ENFORCE(kernelMeta.mD == 32 || kernelMeta.mD == 40 || kernelMeta.mD == 64 || kernelMeta.mD == 80 ||
                kernelMeta.mD == 128 || kernelMeta.mD == 144 || kernelMeta.mD == 160 ||
                kernelMeta.mD == 256);
    return hashID(kernelMeta.mS,
                  kernelMeta.mD,
                  kernelMeta.mInterleaved,
                  kernelMeta.mUnrollStep,
                  kernelMeta.mWithRelativePositionBias,
                  kernelMeta.mFlashAttention,
                  kernelMeta.mCausalMask);
  }

  virtual void run(Fused_multihead_attention_params_v2& params,
                   cudaStream_t ss,
                   bool flash_attention = false,
                   bool causal_mask = false) const {
    ORT_ENFORCE(!params.interleaved);  // only int8 has interleaved

    bool forceUnroll = params.force_unroll;
    if (!forceUnroll && !params.ignore_b1opt && mSM >= kSM_75) {
      const struct {
        unsigned int mSM;
        Data_type mDataType;
        int mS;
        int mD;
        int mMaxBatch;
      } unrollList[] = {
        {kSM_75, DATA_TYPE_FP16, 128, 32, 1},
        {kSM_75, DATA_TYPE_FP16, 128, 40, 1},
        {kSM_75, DATA_TYPE_FP16, 128, 64, 1},
        {kSM_75, DATA_TYPE_FP16, 256, 64, 1},
        {kSM_75, DATA_TYPE_FP16, 384, 64, 1},
#if CUDA_VERSION >= 11000
        {kSM_80, DATA_TYPE_FP16, 128, 32, 4},
        {kSM_80, DATA_TYPE_FP16, 128, 40, 4},
        {kSM_80, DATA_TYPE_FP16, 128, 64, 4},
        {kSM_80, DATA_TYPE_FP16, 256, 64, 4},
        {kSM_80, DATA_TYPE_FP16, 384, 64, 4},
        {kSM_86, DATA_TYPE_FP16, 128, 32, 4},
        {kSM_86, DATA_TYPE_FP16, 128, 40, 4},
        {kSM_86, DATA_TYPE_FP16, 128, 64, 4},
        {kSM_86, DATA_TYPE_FP16, 256, 64, 4},
#endif
      };
      for (unsigned int i = 0u; i < sizeof(unrollList) / sizeof(unrollList[0]); ++i) {
        if (mSM == unrollList[i].mSM && mDataType == unrollList[i].mDataType && params.s == unrollList[i].mS &&
            params.d == unrollList[i].mD && params.b <= unrollList[i].mMaxBatch) {
          forceUnroll = true;
          break;
        }
      }
    }

    bool withRelativePositionBias = params.window_num == 0 ? false : true;
    const auto findIter = mFunctions.find(hashID(params.s,
                                                 params.d,
                                                 params.interleaved,
                                                 forceUnroll,
                                                 withRelativePositionBias,
                                                 flash_attention,
                                                 causal_mask));
    // Provide debug information if the kernel is missing in the pool.
    if (findIter == mFunctions.end()) {
      std::ostringstream errMsg;
      errMsg << "Failed to find kernel: s=" << params.s << " d=" << params.d << " interleaved="
             << params.interleaved << " forceUnroll=" << forceUnroll
             << " withRelativePositionBias=" << withRelativePositionBias
             << " flash_attention=" << flash_attention
             << " causal_mask=" << causal_mask << "\n"
             << "Was the plugin compiled on a compatible CUDA and SM version?\n"
             << "\t Compiled on CUDA " << CUDA_VERSION << "\n"
             << "\t Current SM version: " << mSM << "\n";

      ORT_ENFORCE(findIter != mFunctions.end(), errMsg.str().c_str());
    }

    const auto& kernelMeta = mKernelMeta[findIter->second.mMetaInfoIndex];
    const CUfunction func = findIter->second.mDeviceFunction;
    void* kernelParams[] = {&params, nullptr};
    if (!forceUnroll) {
      cuErrCheck(mDriver.cuLaunchKernel(func,
                                        params.h,
                                        params.b,
                                        1,
                                        kernelMeta.mThreadsPerCTA,
                                        1,
                                        1,
                                        kernelMeta.mSharedMemBytes,
                                        ss,
                                        kernelParams,
                                        nullptr),
                 mDriver);
    } else {
      int unroll = kernelMeta.mS / kernelMeta.mUnrollStep;
      ORT_ENFORCE(kernelMeta.mS == kernelMeta.mUnrollStep * unroll);
      // flash attention supports any sequence length
      if (flash_attention) {
        unroll = (params.s + kernelMeta.mUnrollStep - 1) / kernelMeta.mUnrollStep;
      }
      cuErrCheck(mDriver.cuLaunchKernel(func,
                                        params.h,
                                        params.b,
                                        unroll,
                                        kernelMeta.mThreadsPerCTA,
                                        1,
                                        1,
                                        kernelMeta.mSharedMemBytes,
                                        ss,
                                        kernelParams,
                                        nullptr),
                 mDriver);
    }
  }
};

using FusedMHAKernelFactoryV2 = TFusedMHAKernelFactory<FusedMultiHeadAttentionXMMAKernelV2>;

inline const FusedMultiHeadAttentionXMMAKernelV2* getXMMAKernelsV2(Data_type type, unsigned int sm) {
  return FusedMHAKernelFactoryV2::Get().getXMMAKernels(
      sMhaKernelMetaInfosV2, sizeof(sMhaKernelMetaInfosV2) / sizeof(sMhaKernelMetaInfosV2[0]), type, sm);
}

}  // namespace cuda
}  // namespace contrib
}  // namespace onnxruntime
#if defined(_MSC_VER)
#pragma warning(pop)
#endif